#include "pathgpu/kernels/kernels.h"
std::string luxrays::KernelSource_PathGPU = 
"/***************************************************************************\n"
" *   Copyright (C) 1998-2010 by authors (see AUTHORS.txt )                 *\n"
" *                                                                         *\n"
" *   This file is part of LuxRays.                                         *\n"
" *                                                                         *\n"
" *   LuxRays is free software; you can redistribute it and/or modify       *\n"
" *   it under the terms of the GNU General Public License as published by  *\n"
" *   the Free Software Foundation; either version 3 of the License, or     *\n"
" *   (at your option) any later version.                                   *\n"
" *                                                                         *\n"
" *   LuxRays is distributed in the hope that it will be useful,            *\n"
" *   but WITHOUT ANY WARRANTY; without even the implied warranty of        *\n"
" *   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the         *\n"
" *   GNU General Public License for more details.                          *\n"
" *                                                                         *\n"
" *   You should have received a copy of the GNU General Public License     *\n"
" *   along with this program.  If not, see <http://www.gnu.org/licenses/>. *\n"
" *                                                                         *\n"
" *   LuxRays website: http://www.luxrender.net                             *\n"
" ***************************************************************************/\n"
"\n"
"//#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
"\n"
"// List of symbols defined at compile time:\n"
"//  PARAM_PATH_COUNT\n"
"//  PARAM_IMAGE_WIDTH\n"
"//  PARAM_IMAGE_HEIGHT\n"
"//  PARAM_STARTLINE\n"
"//  PARAM_RASTER2CAMERA_IJ (Matrix4x4)\n"
"//  PARAM_RAY_EPSILON\n"
"//  PARAM_CLIP_YON\n"
"//  PARAM_CLIP_HITHER\n"
"//  PARAM_CAMERA2WORLD_IJ (Matrix4x4)\n"
"//  PARAM_SEED\n"
"//  PARAM_MAX_PATH_DEPTH\n"
"//  PARAM_MAX_RR_DEPTH\n"
"//  PARAM_MAX_RR_CAP\n"
"//  PARAM_SAMPLE_PER_PIXEL\n"
"//  PARAM_CAMERA_HAS_DOF\n"
"//  PARAM_CAMERA_LENS_RADIUS\n"
"//  PARAM_CAMERA_FOCAL_DISTANCE\n"
"//  PARAM_OPENGL_INTEROP\n"
"//  PARAM_DIRECT_LIGHT_SAMPLING\n"
"//  PARAM_DL_LIGHT_COUNT\n"
"//  PARAM_CAMERA_DYNAMIC\n"
"//  PARAM_HAS_TEXTUREMAPS\n"
"//  PARAM_HAS_ALPHA_TEXTUREMAPS\n"
"\n"
"// To enable single material suopport (work around for ATI compiler problems)\n"
"//  PARAM_ENABLE_MAT_MATTE\n"
"//  PARAM_ENABLE_MAT_AREALIGHT\n"
"//  PARAM_ENABLE_MAT_MIRROR\n"
"//  PARAM_ENABLE_MAT_GLASS\n"
"//  PARAM_ENABLE_MAT_MATTEMIRROR\n"
"//  PARAM_ENABLE_MAT_METAL\n"
"//  PARAM_ENABLE_MAT_MATTEMETAL\n"
"//  PARAM_ENABLE_MAT_ALLOY\n"
"//  PARAM_ENABLE_MAT_ARCHGLASS\n"
"\n"
"// (optional)\n"
"//  PARAM_HAVE_INFINITELIGHT\n"
"//  PARAM_IL_GAIN_R\n"
"//  PARAM_IL_GAIN_G\n"
"//  PARAM_IL_GAIN_B\n"
"//  PARAM_IL_SHIFT_U\n"
"//  PARAM_IL_SHIFT_V\n"
"//  PARAM_IL_WIDTH\n"
"//  PARAM_IL_HEIGHT\n"
"\n"
"#ifndef M_PI\n"
"#define M_PI 3.14159265358979323846f\n"
"#endif\n"
"\n"
"#ifndef INV_PI\n"
"#define INV_PI  0.31830988618379067154f\n"
"#endif\n"
"\n"
"#ifndef INV_TWOPI\n"
"#define INV_TWOPI  0.15915494309189533577f\n"
"#endif\n"
"\n"
"#ifndef TRUE\n"
"#define TRUE 1\n"
"#endif\n"
"\n"
"#ifndef FALSE\n"
"#define FALSE 0\n"
"#endif\n"
"\n"
"//------------------------------------------------------------------------------\n"
"// Types\n"
"//------------------------------------------------------------------------------\n"
"\n"
"typedef struct {\n"
"	float u, v;\n"
"} UV;\n"
"\n"
"typedef struct {\n"
"	float r, g, b;\n"
"} Spectrum;\n"
"\n"
"typedef struct {\n"
"	float x, y, z;\n"
"} Point;\n"
"\n"
"typedef struct {\n"
"	float x, y, z;\n"
"} Vector;\n"
"\n"
"typedef struct {\n"
"	uint v0, v1, v2;\n"
"} Triangle;\n"
"\n"
"typedef struct {\n"
"	Point o;\n"
"	Vector d;\n"
"	float mint, maxt;\n"
"} Ray;\n"
"\n"
"typedef struct {\n"
"	float t;\n"
"	float b1, b2; // Barycentric coordinates of the hit point\n"
"	uint index;\n"
"} RayHit;\n"
"\n"
"typedef struct {\n"
"	uint s1, s2, s3;\n"
"} Seed;\n"
"\n"
"#define PATH_STATE_NEXT_VERTEX 0\n"
"#define PATH_STATE_SAMPLE_LIGHT 1\n"
"#define PATH_STATE_CONTINUE_NEXT_VERTEX 2\n"
"\n"
"typedef struct {\n"
"	Spectrum throughput;\n"
"	uint depth, pixelIndex, subpixelIndex;\n"
"	Seed seed;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"	int specularBounce;\n"
"	int state;\n"
"\n"
"	Ray pathRay;\n"
"	RayHit pathHit;\n"
"	Spectrum lightRadiance;\n"
"	Spectrum accumRadiance;\n"
"#endif\n"
"} Path;\n"
"\n"
"typedef struct {\n"
"	Spectrum c;\n"
"	uint count;\n"
"} Pixel;\n"
"\n"
"#define MAT_MATTE 0\n"
"#define MAT_AREALIGHT 1\n"
"#define MAT_MIRROR 2\n"
"#define MAT_GLASS 3\n"
"#define MAT_MATTEMIRROR 4\n"
"#define MAT_METAL 5\n"
"#define MAT_MATTEMETAL 6\n"
"#define MAT_ALLOY 7\n"
"#define MAT_ARCHGLASS 8\n"
"#define MAT_NULL 9\n"
"\n"
"typedef struct {\n"
"    float r, g, b;\n"
"} MatteParam;\n"
"\n"
"typedef struct {\n"
"    float gain_r, gain_g, gain_b;\n"
"} AreaLightParam;\n"
"\n"
"typedef struct {\n"
"    float r, g, b;\n"
"    int specularBounce;\n"
"} MirrorParam;\n"
"\n"
"typedef struct {\n"
"    float refl_r, refl_g, refl_b;\n"
"    float refrct_r, refrct_g, refrct_b;\n"
"    float ousideIor, ior;\n"
"    float R0;\n"
"    int reflectionSpecularBounce, transmitionSpecularBounce;\n"
"} GlassParam;\n"
"\n"
"typedef struct {\n"
"	MatteParam matte;\n"
"	MirrorParam mirror;\n"
"	float matteFilter, totFilter, mattePdf, mirrorPdf;\n"
"} MatteMirrorParam;\n"
"\n"
"typedef struct {\n"
"    float r, g, b;\n"
"    float exponent;\n"
"    int specularBounce;\n"
"} MetalParam;\n"
"\n"
"typedef struct {\n"
"	MatteParam matte;\n"
"	MetalParam metal;\n"
"	float matteFilter, totFilter, mattePdf, metalPdf;\n"
"} MatteMetalParam;\n"
"\n"
"typedef struct {\n"
"    float diff_r, diff_g, diff_b;\n"
"    float refl_r, refl_g, refl_b;\n"
"    float exponent;\n"
"    float R0;\n"
"    int specularBounce;\n"
"} AlloyParam;\n"
"\n"
"typedef struct {\n"
"    float refl_r, refl_g, refl_b;\n"
"    float refrct_r, refrct_g, refrct_b;\n"
"	float transFilter, totFilter, reflPdf, transPdf;\n"
"	bool reflectionSpecularBounce, transmitionSpecularBounce;\n"
"} ArchGlassParam;\n"
"\n"
"typedef struct {\n"
"	unsigned int type;\n"
"	union {\n"
"		MatteParam matte;\n"
"        AreaLightParam areaLight;\n"
"		MirrorParam mirror;\n"
"        GlassParam glass;\n"
"		MatteMirrorParam matteMirror;\n"
"        MetalParam metal;\n"
"        MatteMetalParam matteMetal;\n"
"        AlloyParam alloy;\n"
"        ArchGlassParam archGlass;\n"
"	} param;\n"
"} Material;\n"
"\n"
"typedef struct {\n"
"	Point v0, v1, v2;\n"
"	Vector normal;\n"
"	float area;\n"
"	float gain_r, gain_g, gain_b;\n"
"} TriangleLight;\n"
"\n"
"typedef struct {\n"
"	unsigned int rgbOffset, alphaOffset;\n"
"	unsigned int width, height;\n"
"} TexMap;\n"
"\n"
"//------------------------------------------------------------------------------\n"
"// Random number generator\n"
"// maximally equidistributed combined Tausworthe generator\n"
"//------------------------------------------------------------------------------\n"
"\n"
"#define FLOATMASK 0x00ffffffu\n"
"\n"
"uint TAUSWORTHE(const uint s, const uint a,\n"
"	const uint b, const uint c,\n"
"	const uint d) {\n"
"	return ((s&c)<<d) ^ (((s << a) ^ s) >> b);\n"
"}\n"
"\n"
"uint LCG(const uint x) { return x * 69069; }\n"
"\n"
"uint ValidSeed(const uint x, const uint m) {\n"
"	return (x < m) ? (x + m) : x;\n"
"}\n"
"\n"
"void InitRandomGenerator(uint seed, Seed *s) {\n"
"	// Avoid 0 value\n"
"	seed = (seed == 0) ? (seed + 0xffffffu) : seed;\n"
"\n"
"	s->s1 = ValidSeed(LCG(seed), 1);\n"
"	s->s2 = ValidSeed(LCG(s->s1), 7);\n"
"	s->s3 = ValidSeed(LCG(s->s2), 15);\n"
"}\n"
"\n"
"unsigned long RndUintValue(Seed *s) {\n"
"	s->s1 = TAUSWORTHE(s->s1, 13, 19, 4294967294UL, 12);\n"
"	s->s2 = TAUSWORTHE(s->s2, 2, 25, 4294967288UL, 4);\n"
"	s->s3 = TAUSWORTHE(s->s3, 3, 11, 4294967280UL, 17);\n"
"\n"
"	return ((s->s1) ^ (s->s2) ^ (s->s3));\n"
"}\n"
"\n"
"float RndFloatValue(Seed *s) {\n"
"	return (RndUintValue(s) & FLOATMASK) * (1.f / (FLOATMASK + 1UL));\n"
"}\n"
"\n"
"//------------------------------------------------------------------------------\n"
"\n"
"float Dot(const Vector *v0, const Vector *v1) {\n"
"	return v0->x * v1->x + v0->y * v1->y + v0->z * v1->z;\n"
"}\n"
"\n"
"void Normalize(Vector *v) {\n"
"	const float il = 1.f / sqrt(Dot(v, v));\n"
"\n"
"	v->x *= il;\n"
"	v->y *= il;\n"
"	v->z *= il;\n"
"}\n"
"\n"
"void Cross(Vector *v3, const Vector *v1, const Vector *v2) {\n"
"	v3->x = (v1->y * v2->z) - (v1->z * v2->y);\n"
"	v3->y = (v1->z * v2->x) - (v1->x * v2->z),\n"
"	v3->z = (v1->x * v2->y) - (v1->y * v2->x);\n"
"}\n"
"\n"
"void ConcentricSampleDisk(const float u1, const float u2, float *dx, float *dy) {\n"
"	float r, theta;\n"
"	// Map uniform random numbers to $[-1,1]^2$\n"
"	float sx = 2.f * u1 - 1.f;\n"
"	float sy = 2.f * u2 - 1.f;\n"
"	// Map square to $(r,\theta)$\n"
"	// Handle degeneracy at the origin\n"
"	if (sx == 0.f && sy == 0.f) {\n"
"		*dx = 0.f;\n"
"		*dy = 0.f;\n"
"		return;\n"
"	}\n"
"	if (sx >= -sy) {\n"
"		if (sx > sy) {\n"
"			// Handle first region of disk\n"
"			r = sx;\n"
"			if (sy > 0.f)\n"
"				theta = sy / r;\n"
"			else\n"
"				theta = 8.f + sy / r;\n"
"		} else {\n"
"			// Handle second region of disk\n"
"			r = sy;\n"
"			theta = 2.f - sx / r;\n"
"		}\n"
"	} else {\n"
"		if (sx <= sy) {\n"
"			// Handle third region of disk\n"
"			r = -sx;\n"
"			theta = 4.f - sy / r;\n"
"		} else {\n"
"			// Handle fourth region of disk\n"
"			r = -sy;\n"
"			theta = 6.f + sx / r;\n"
"		}\n"
"	}\n"
"	theta *= M_PI / 4.f;\n"
"	*dx = r * cos(theta);\n"
"	*dy = r * sin(theta);\n"
"}\n"
"\n"
"void CosineSampleHemisphere(Vector *ret, const float u1, const float u2) {\n"
"	ConcentricSampleDisk(u1, u2, &ret->x, &ret->y);\n"
"	ret->z = sqrt(max(0.f, 1.f - ret->x * ret->x - ret->y * ret->y));\n"
"}\n"
"\n"
"void CoordinateSystem(const Vector *v1, Vector *v2, Vector *v3) {\n"
"	if (fabs(v1->x) > fabs(v1->y)) {\n"
"		float invLen = 1.f / sqrt(v1->x * v1->x + v1->z * v1->z);\n"
"		v2->x = -v1->z * invLen;\n"
"		v2->y = 0.f;\n"
"		v2->z = v1->x * invLen;\n"
"	} else {\n"
"		float invLen = 1.f / sqrt(v1->y * v1->y + v1->z * v1->z);\n"
"		v2->x = 0.f;\n"
"		v2->y = v1->z * invLen;\n"
"		v2->z = -v1->y * invLen;\n"
"	}\n"
"\n"
"	Cross(v3, v1, v2);\n"
"}\n"
"\n"
"//------------------------------------------------------------------------------\n"
"\n"
"void GenerateRay(\n"
"		const uint pixelIndex,\n"
"		__global Ray *ray, Seed *seed\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"		, __global float *cameraData\n"
"#endif\n"
"		) {\n"
"\n"
"	/*// Gaussina distribution\n"
"	const float rad = filterRadius * sqrt(-log(1.f - RndFloatValue(seed)));\n"
"	const float angle = 2 * M_PI * RndFloatValue(seed);\n"
"\n"
"	const float screenX = pixelIndex % PARAM_IMAGE_WIDTH + rad * cos(angle);\n"
"	const float screenY = pixelIndex / PARAM_IMAGE_WIDTH + rad * sin(angle);*/\n"
"\n"
"	const float screenX = pixelIndex % PARAM_IMAGE_WIDTH + RndFloatValue(seed) - 0.5f;\n"
"	const float screenY = pixelIndex / PARAM_IMAGE_WIDTH + RndFloatValue(seed) - 0.5f;\n"
"\n"
"	Point Pras;\n"
"	Pras.x = screenX;\n"
"	Pras.y = PARAM_IMAGE_HEIGHT - screenY - 1.f;\n"
"	Pras.z = 0;\n"
"\n"
"	Point orig;\n"
"	// RasterToCamera(Pras, &orig);\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"	const float iw = 1.f / (cameraData[12] * Pras.x + cameraData[13] * Pras.y + cameraData[14] * Pras.z + cameraData[15]);\n"
"	orig.x = (cameraData[0] * Pras.x + cameraData[1] * Pras.y + cameraData[2] * Pras.z + cameraData[3]) * iw;\n"
"	orig.y = (cameraData[4] * Pras.x + cameraData[5] * Pras.y + cameraData[6] * Pras.z + cameraData[7]) * iw;\n"
"	orig.z = (cameraData[8] * Pras.x + cameraData[9] * Pras.y + cameraData[10] * Pras.z + cameraData[11]) * iw;\n"
"#else\n"
"	const float iw = 1.f / (PARAM_RASTER2CAMERA_30 * Pras.x + PARAM_RASTER2CAMERA_31 * Pras.y + PARAM_RASTER2CAMERA_32 * Pras.z + PARAM_RASTER2CAMERA_33);\n"
"	orig.x = (PARAM_RASTER2CAMERA_00 * Pras.x + PARAM_RASTER2CAMERA_01 * Pras.y + PARAM_RASTER2CAMERA_02 * Pras.z + PARAM_RASTER2CAMERA_03) * iw;\n"
"	orig.y = (PARAM_RASTER2CAMERA_10 * Pras.x + PARAM_RASTER2CAMERA_11 * Pras.y + PARAM_RASTER2CAMERA_12 * Pras.z + PARAM_RASTER2CAMERA_13) * iw;\n"
"	orig.z = (PARAM_RASTER2CAMERA_20 * Pras.x + PARAM_RASTER2CAMERA_21 * Pras.y + PARAM_RASTER2CAMERA_22 * Pras.z + PARAM_RASTER2CAMERA_23) * iw;\n"
"#endif\n"
"\n"
"	Vector dir;\n"
"	dir.x = orig.x;\n"
"	dir.y = orig.y;\n"
"	dir.z = orig.z;\n"
"\n"
"#if defined(PARAM_CAMERA_HAS_DOF)\n"
"	// Sample point on lens\n"
"	float lensU, lensV;\n"
"	ConcentricSampleDisk(RndFloatValue(seed), RndFloatValue(seed), &lensU, &lensV);\n"
"	lensU *= PARAM_CAMERA_LENS_RADIUS;\n"
"	lensV *= PARAM_CAMERA_LENS_RADIUS;\n"
"\n"
"	// Compute point on plane of focus\n"
"	const float ft = (PARAM_CAMERA_FOCAL_DISTANCE - PARAM_CLIP_HITHER) / dir.z;\n"
"	Point Pfocus;\n"
"	Pfocus.x = orig.x + dir.x * ft;\n"
"	Pfocus.y = orig.y + dir.y * ft;\n"
"	Pfocus.z = orig.z + dir.z * ft;\n"
"\n"
"	// Update ray for effect of lens\n"
"	orig.x += lensU * ((PARAM_CAMERA_FOCAL_DISTANCE - PARAM_CLIP_HITHER) / PARAM_CAMERA_FOCAL_DISTANCE);\n"
"	orig.y += lensV * ((PARAM_CAMERA_FOCAL_DISTANCE - PARAM_CLIP_HITHER) / PARAM_CAMERA_FOCAL_DISTANCE);\n"
"\n"
"	dir.x = Pfocus.x - orig.x;\n"
"	dir.y = Pfocus.y - orig.y;\n"
"	dir.z = Pfocus.z - orig.z;\n"
"#endif\n"
"\n"
"	Normalize(&dir);\n"
"\n"
"	// CameraToWorld(*ray, ray);\n"
"	Point torig;\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"	const float iw2 = 1.f / (cameraData[16 + 12] * orig.x + cameraData[16 + 13] * orig.y + cameraData[16 + 14] * orig.z + cameraData[16 + 15]);\n"
"	torig.x = (cameraData[16 + 0] * orig.x + cameraData[16 + 1] * orig.y + cameraData[16 + 2] * orig.z + cameraData[16 + 3]) * iw2;\n"
"	torig.y = (cameraData[16 + 4] * orig.x + cameraData[16 + 5] * orig.y + cameraData[16 + 6] * orig.z + cameraData[16 + 7]) * iw2;\n"
"	torig.z = (cameraData[16 + 8] * orig.x + cameraData[16 + 9] * orig.y + cameraData[16 + 12] * orig.z + cameraData[16 + 11]) * iw2;\n"
"#else\n"
"	const float iw2 = 1.f / (PARAM_CAMERA2WORLD_30 * orig.x + PARAM_CAMERA2WORLD_31 * orig.y + PARAM_CAMERA2WORLD_32 * orig.z + PARAM_CAMERA2WORLD_33);\n"
"	torig.x = (PARAM_CAMERA2WORLD_00 * orig.x + PARAM_CAMERA2WORLD_01 * orig.y + PARAM_CAMERA2WORLD_02 * orig.z + PARAM_CAMERA2WORLD_03) * iw2;\n"
"	torig.y = (PARAM_CAMERA2WORLD_10 * orig.x + PARAM_CAMERA2WORLD_11 * orig.y + PARAM_CAMERA2WORLD_12 * orig.z + PARAM_CAMERA2WORLD_13) * iw2;\n"
"	torig.z = (PARAM_CAMERA2WORLD_20 * orig.x + PARAM_CAMERA2WORLD_21 * orig.y + PARAM_CAMERA2WORLD_22 * orig.z + PARAM_CAMERA2WORLD_23) * iw2;\n"
"#endif\n"
"\n"
"	Vector tdir;\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"	tdir.x = cameraData[16 + 0] * dir.x + cameraData[16 + 1] * dir.y + cameraData[16 + 2] * dir.z;\n"
"	tdir.y = cameraData[16 + 4] * dir.x + cameraData[16 + 5] * dir.y + cameraData[16 + 6] * dir.z;\n"
"	tdir.z = cameraData[16 + 8] * dir.x + cameraData[16 + 9] * dir.y + cameraData[16 + 10] * dir.z;\n"
"#else\n"
"	tdir.x = PARAM_CAMERA2WORLD_00 * dir.x + PARAM_CAMERA2WORLD_01 * dir.y + PARAM_CAMERA2WORLD_02 * dir.z;\n"
"	tdir.y = PARAM_CAMERA2WORLD_10 * dir.x + PARAM_CAMERA2WORLD_11 * dir.y + PARAM_CAMERA2WORLD_12 * dir.z;\n"
"	tdir.z = PARAM_CAMERA2WORLD_20 * dir.x + PARAM_CAMERA2WORLD_21 * dir.y + PARAM_CAMERA2WORLD_22 * dir.z;\n"
"#endif\n"
"\n"
"	ray->o = torig;\n"
"	ray->d = tdir;\n"
"	ray->mint = PARAM_RAY_EPSILON;\n"
"	ray->maxt = (PARAM_CLIP_YON - PARAM_CLIP_HITHER) / dir.z;\n"
"}\n"
"\n"
"//------------------------------------------------------------------------------\n"
"\n"
"__kernel void Init(\n"
"		__global Path *paths,\n"
"		__global Ray *rays\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"		, __global float *cameraData\n"
"#endif\n"
"		) {\n"
"	const int gid = get_global_id(0);\n"
"	if (gid >= PARAM_PATH_COUNT)\n"
"		return;\n"
"\n"
"	// Initialize the path\n"
"	__global Path *path = &paths[gid];\n"
"	path->throughput.r = 1.f;\n"
"	path->throughput.g = 1.f;\n"
"	path->throughput.b = 1.f;\n"
"	path->depth = 0;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"	path->specularBounce = TRUE;\n"
"	path->state = PATH_STATE_NEXT_VERTEX;\n"
"	path->accumRadiance.r = 0.f;\n"
"	path->accumRadiance.g = 0.f;\n"
"	path->accumRadiance.b = 0.f;\n"
"#endif\n"
"\n"
"	const uint pixelIndex = (PARAM_STARTLINE * PARAM_IMAGE_WIDTH + gid) % (PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT);\n"
"	path->pixelIndex = pixelIndex;\n"
"	path->subpixelIndex = 0;\n"
"\n"
"	// Initialize random number generator\n"
"	Seed seed;\n"
"	InitRandomGenerator(PARAM_SEED + gid, &seed);\n"
"\n"
"	// Generate the eye ray\n"
"	GenerateRay(pixelIndex, &rays[gid], &seed\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"		, cameraData\n"
"#endif\n"
"		);\n"
"\n"
"	// Save the seed\n"
"	path->seed.s1 = seed.s1;\n"
"	path->seed.s2 = seed.s2;\n"
"	path->seed.s3 = seed.s3;\n"
"}\n"
"\n"
"//------------------------------------------------------------------------------\n"
"\n"
"__kernel void InitFrameBuffer(\n"
"		__global Pixel *frameBuffer\n"
"#if defined(PARAM_OPENGL_INTEROP)\n"
"		, const int clearPBO\n"
"		, __global uint *pbo\n"
"#endif\n"
"		) {\n"
"	const int gid = get_global_id(0);\n"
"	if (gid >= PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT)\n"
"		return;\n"
"\n"
"	__global Pixel *p = &frameBuffer[gid];\n"
"	p->c.r = 0.f;\n"
"	p->c.g = 0.f;\n"
"	p->c.b = 0.f;\n"
"	p->count = 0;\n"
"\n"
"#if defined(PARAM_OPENGL_INTEROP)\n"
"	if (clearPBO)\n"
"		pbo[gid] = 0;\n"
"#endif\n"
"}\n"
"\n"
"#if defined(PARAM_OPENGL_INTEROP)\n"
"uint Radiance2PixelUInt(const float x) {\n"
"	return (uint)(pow(clamp(x, 0.f, 1.f), 1.f / 2.2f) * 255.f + .5f);\n"
"}\n"
"\n"
"__kernel void UpdatePixelBuffer(\n"
"		__global Pixel *frameBuffer,\n"
"		__global uint *pbo) {\n"
"	const int gid = get_global_id(0);\n"
"	if (gid >= PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT)\n"
"		return;\n"
"\n"
"	__global Pixel *p = &frameBuffer[gid];\n"
"\n"
"	const uint count = p->count;\n"
"	if (count > 0) {\n"
"		const float invCount = 1.f / p->count;\n"
"\n"
"		const uint r = Radiance2PixelUInt(p->c.r * invCount);\n"
"		const uint g = Radiance2PixelUInt(p->c.g * invCount);\n"
"		const uint b = Radiance2PixelUInt(p->c.b * invCount);\n"
"		pbo[gid] = r | (g << 8) | (b << 16);\n"
"	}\n"
"}\n"
"\n"
"void AccumPixel(__global Pixel *p, float *r, float *g, float *b, const float k) {\n"
"	const uint count = p->count;\n"
"\n"
"	if (count > 0) {\n"
"		const float invCount = k / p->count;\n"
"\n"
"		*r += p->c.r * invCount;\n"
"		*g += p->c.g * invCount;\n"
"		*b += p->c.b * invCount;\n"
"	}\n"
"}\n"
"\n"
"__kernel void UpdatePixelBufferBlured(\n"
"		__global Pixel *frameBuffer,\n"
"		__global uint *pbo) {\n"
"	const int gid = get_global_id(0);\n"
"	if (gid >= PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT)\n"
"		return;\n"
"\n"
"	__global Pixel *p = &frameBuffer[gid];\n"
"\n"
"	const uint count = p->count;\n"
"	if (count > 0) {\n"
"		// Naive Gaussian Filter (i.e. not using local memory, etc.)\n"
"\n"
"		const unsigned int x = gid % PARAM_IMAGE_WIDTH;\n"
"		const unsigned int y = gid / PARAM_IMAGE_WIDTH;\n"
"\n"
"		float r = 0.f;\n"
"		float g = 0.f;\n"
"		float b = 0.f;\n"
"\n"
"		if (y > 0) {\n"
"			if (x > 0) AccumPixel(&frameBuffer[gid - PARAM_IMAGE_WIDTH - 1], &r, &g, &b, 1.f / 16.f);\n"
"			AccumPixel(&frameBuffer[gid - PARAM_IMAGE_WIDTH], &r, &g, &b, 2.f / 16.f);\n"
"			if (x < PARAM_IMAGE_WIDTH - 1) AccumPixel(&frameBuffer[gid - PARAM_IMAGE_WIDTH + 1], &r, &g, &b, 1.f / 16.f);\n"
"		}\n"
"\n"
"		if (x > 0) AccumPixel(&frameBuffer[gid - 1], &r, &g, &b, 2.f / 16.f);\n"
"		AccumPixel(&frameBuffer[gid], &r, &g, &b, 4.f / 16.f);\n"
"		if (x < PARAM_IMAGE_WIDTH - 1) AccumPixel(&frameBuffer[gid + 1], &r, &g, &b, 2.f / 16.f);\n"
"\n"
"		if (y < PARAM_IMAGE_HEIGHT - 1) {\n"
"			if (x > 0) AccumPixel(&frameBuffer[gid + PARAM_IMAGE_WIDTH - 1], &r, &g, &b, 1.f / 16.f);\n"
"			AccumPixel(&frameBuffer[gid + PARAM_IMAGE_WIDTH], &r, &g, &b, 2.f / 16.f);\n"
"			if (x < PARAM_IMAGE_WIDTH - 1) AccumPixel(&frameBuffer[gid + PARAM_IMAGE_WIDTH + 1], &r, &g, &b, 1.f / 16.f);\n"
"		}\n"
"\n"
"		const uint pr = Radiance2PixelUInt(r);\n"
"		const uint pg = Radiance2PixelUInt(g);\n"
"		const uint pb = Radiance2PixelUInt(b);\n"
"		pbo[gid] = pr | (pg << 8) | (pb << 16);\n"
"	}\n"
"}\n"
"\n"
"#endif\n"
"\n"
"//------------------------------------------------------------------------------\n"
"\n"
"int Mod(int a, int b) {\n"
"	if (b == 0)\n"
"		b = 1;\n"
"\n"
"	a %= b;\n"
"	if (a < 0)\n"
"		a += b;\n"
"\n"
"	return a;\n"
"}\n"
"\n"
"void TexMap_GetTexel(__global Spectrum *pixels, const uint width, const uint height,\n"
"		const int s, const int t, Spectrum *col) {\n"
"	const uint u = Mod(s, width);\n"
"	const uint v = Mod(t, height);\n"
"\n"
"	const unsigned index = v * width + u;\n"
"\n"
"	col->r = pixels[index].r;\n"
"	col->g = pixels[index].g;\n"
"	col->b = pixels[index].b;\n"
"}\n"
"\n"
"float TexMap_GetAlphaTexel(__global float *alphas, const uint width, const uint height,\n"
"		const int s, const int t) {\n"
"	const uint u = Mod(s, width);\n"
"	const uint v = Mod(t, height);\n"
"\n"
"	const unsigned index = v * width + u;\n"
"\n"
"	return alphas[index];\n"
"}\n"
"\n"
"void TexMap_GetColor(__global Spectrum *pixels, const uint width, const uint height,\n"
"		const float u, const float v, Spectrum *col) {\n"
"	const float s = u * width - 0.5f;\n"
"	const float t = v * height - 0.5f;\n"
"\n"
"	const int s0 = (int)floor(s);\n"
"	const int t0 = (int)floor(t);\n"
"\n"
"	const float ds = s - s0;\n"
"	const float dt = t - t0;\n"
"\n"
"	const float ids = 1.f - ds;\n"
"	const float idt = 1.f - dt;\n"
"\n"
"	Spectrum c0, c1, c2, c3;\n"
"	TexMap_GetTexel(pixels, width, height, s0, t0, &c0);\n"
"	TexMap_GetTexel(pixels, width, height, s0, t0 + 1, &c1);\n"
"	TexMap_GetTexel(pixels, width, height, s0 + 1, t0, &c2);\n"
"	TexMap_GetTexel(pixels, width, height, s0 + 1, t0 + 1, &c3);\n"
"\n"
"	const float k0 = ids * idt;\n"
"	const float k1 = ids * dt;\n"
"	const float k2 = ds * idt;\n"
"	const float k3 = ds * dt;\n"
"\n"
"	col->r = k0 * c0.r + k1 * c1.r + k2 * c2.r + k3 * c3.r;\n"
"	col->g = k0 * c0.g + k1 * c1.g + k2 * c2.g + k3 * c3.g;\n"
"	col->b = k0 * c0.b + k1 * c1.b + k2 * c2.b + k3 * c3.b;\n"
"}\n"
"\n"
"float TexMap_GetAlpha(__global float *alphas, const uint width, const uint height,\n"
"		const float u, const float v) {\n"
"	const float s = u * width - 0.5f;\n"
"	const float t = v * height - 0.5f;\n"
"\n"
"	const int s0 = (int)floor(s);\n"
"	const int t0 = (int)floor(t);\n"
"\n"
"	const float ds = s - s0;\n"
"	const float dt = t - t0;\n"
"\n"
"	const float ids = 1.f - ds;\n"
"	const float idt = 1.f - dt;\n"
"\n"
"	const float c0 = TexMap_GetAlphaTexel(alphas, width, height, s0, t0);\n"
"	const float c1 = TexMap_GetAlphaTexel(alphas, width, height, s0, t0 + 1);\n"
"	const float c2 = TexMap_GetAlphaTexel(alphas, width, height, s0 + 1, t0);\n"
"	const float c3 = TexMap_GetAlphaTexel(alphas, width, height, s0 + 1, t0 + 1);\n"
"\n"
"	const float k0 = ids * idt;\n"
"	const float k1 = ids * dt;\n"
"	const float k2 = ds * idt;\n"
"	const float k3 = ds * dt;\n"
"\n"
"	return k0 * c0 + k1 * c1 + k2 * c2 + k3 * c3;\n"
"}\n"
"\n"
"float SphericalTheta(const Vector *v) {\n"
"	return acos(clamp(v->z, -1.f, 1.f));\n"
"}\n"
"\n"
"float SphericalPhi(const Vector *v) {\n"
"	float p = atan2(v->y, v->x);\n"
"	return (p < 0.f) ? p + 2.f * M_PI : p;\n"
"}\n"
"\n"
"#if defined(PARAM_HAVE_INFINITELIGHT)\n"
"void InfiniteLight_Le(__global Spectrum *infiniteLightMap, Spectrum *le, Vector *dir) {\n"
"	const float u = 1.f - SphericalPhi(dir) * INV_TWOPI +  PARAM_IL_SHIFT_U;\n"
"	const float v = SphericalTheta(dir) * INV_PI + PARAM_IL_SHIFT_V;\n"
"\n"
"	TexMap_GetColor(infiniteLightMap, PARAM_IL_WIDTH, PARAM_IL_HEIGHT, u, v, le);\n"
"\n"
"	le->r *= PARAM_IL_GAIN_R;\n"
"	le->g *= PARAM_IL_GAIN_G;\n"
"	le->b *= PARAM_IL_GAIN_B;\n"
"}\n"
"#endif\n"
"\n"
"void Mesh_InterpolateColor(__global Spectrum *colors, __global Triangle *triangles,\n"
"		const uint triIndex, const float b1, const float b2, Spectrum *C) {\n"
"	__global Triangle *tri = &triangles[triIndex];\n"
"\n"
"	const float b0 = 1.f - b1 - b2;\n"
"	C->r = b0 * colors[tri->v0].r + b1 * colors[tri->v1].r + b2 * colors[tri->v2].r;\n"
"	C->g = b0 * colors[tri->v0].g + b1 * colors[tri->v1].g + b2 * colors[tri->v2].g;\n"
"	C->b = b0 * colors[tri->v0].b + b1 * colors[tri->v1].b + b2 * colors[tri->v2].b;\n"
"}\n"
"\n"
"void Mesh_InterpolateNormal(__global Vector *normals, __global Triangle *triangles,\n"
"		const uint triIndex, const float b1, const float b2, Vector *N) {\n"
"	__global Triangle *tri = &triangles[triIndex];\n"
"\n"
"	const float b0 = 1.f - b1 - b2;\n"
"	N->x = b0 * normals[tri->v0].x + b1 * normals[tri->v1].x + b2 * normals[tri->v2].x;\n"
"	N->y = b0 * normals[tri->v0].y + b1 * normals[tri->v1].y + b2 * normals[tri->v2].y;\n"
"	N->z = b0 * normals[tri->v0].z + b1 * normals[tri->v1].z + b2 * normals[tri->v2].z;\n"
"\n"
"	Normalize(N);\n"
"}\n"
"\n"
"void Mesh_InterpolateUV(__global UV *uvs, __global Triangle *triangles,\n"
"		const uint triIndex, const float b1, const float b2, UV *uv) {\n"
"	__global Triangle *tri = &triangles[triIndex];\n"
"\n"
"	const float b0 = 1.f - b1 - b2;\n"
"	uv->u = b0 * uvs[tri->v0].u + b1 * uvs[tri->v1].u + b2 * uvs[tri->v2].u;\n"
"	uv->v = b0 * uvs[tri->v0].v + b1 * uvs[tri->v1].v + b2 * uvs[tri->v2].v;\n"
"}\n"
"\n"
"//------------------------------------------------------------------------------\n"
"// Materials\n"
"//------------------------------------------------------------------------------\n"
"\n"
"void Matte_Sample_f(__global MatteParam *mat, const Vector *wo, Vector *wi,\n"
"		float *pdf, Spectrum *f, const Vector *shadeN,\n"
"		const float u0, const float u1\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"		) {\n"
"	Vector dir;\n"
"	CosineSampleHemisphere(&dir, u0, u1);\n"
"	*pdf = dir.z * INV_PI;\n"
"\n"
"	Vector v1, v2;\n"
"	CoordinateSystem(shadeN, &v1, &v2);\n"
"\n"
"	wi->x = v1.x * dir.x + v2.x * dir.y + shadeN->x * dir.z;\n"
"	wi->y = v1.y * dir.x + v2.y * dir.y + shadeN->y * dir.z;\n"
"	wi->z = v1.z * dir.x + v2.z * dir.y + shadeN->z * dir.z;\n"
"\n"
"	const float dp = Dot(shadeN, wi);\n"
"	// Using 0.0001 instead of 0.0 to cut down fireflies\n"
"	if (dp <= 0.0001f)\n"
"		*pdf = 0.f;\n"
"	else {\n"
"		*pdf /=  dp;\n"
"\n"
"		f->r = mat->r * INV_PI;\n"
"		f->g = mat->g * INV_PI;\n"
"		f->b = mat->b * INV_PI;\n"
"	}\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"	*specularBounce = FALSE;\n"
"#endif\n"
"}\n"
"\n"
"void Mirror_Sample_f(__global MirrorParam *mat, const Vector *wo, Vector *wi,\n"
"		float *pdf, Spectrum *f, const Vector *shadeN\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"		) {\n"
"    const float k = 2.f * Dot(shadeN, wo);\n"
"	wi->x = k * shadeN->x - wo->x;\n"
"	wi->y = k * shadeN->y - wo->y;\n"
"	wi->z = k * shadeN->z - wo->z;\n"
"\n"
"	*pdf = 1.f;\n"
"\n"
"	f->r = mat->r;\n"
"	f->g = mat->g;\n"
"	f->b = mat->b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"	*specularBounce = mat->specularBounce;\n"
"#endif\n"
"}\n"
"\n"
"void Glass_Sample_f(__global GlassParam *mat,\n"
"    const Vector *wo, Vector *wi, float *pdf, Spectrum *f, const Vector *N, const Vector *shadeN,\n"
"    const float u0\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"        ) {\n"
"    Vector reflDir;\n"
"    const float k = 2.f * Dot(N, wo);\n"
"    reflDir.x = k * N->x - wo->x;\n"
"    reflDir.y = k * N->y - wo->y;\n"
"    reflDir.z = k * N->z - wo->z;\n"
"\n"
"    // Ray from outside going in ?\n"
"    const bool into = (Dot(N, shadeN) > 0.f);\n"
"\n"
"    const float nc = mat->ousideIor;\n"
"    const float nt = mat->ior;\n"
"    const float nnt = into ? (nc / nt) : (nt / nc);\n"
"    const float ddn = -Dot(wo, shadeN);\n"
"    const float cos2t = 1.f - nnt * nnt * (1.f - ddn * ddn);\n"
"\n"
"    // Total internal reflection\n"
"    if (cos2t < 0.f) {\n"
"        *wi = reflDir;\n"
"        *pdf = 1.f;\n"
"\n"
"        f->r = mat->refl_r;\n"
"        f->g = mat->refl_g;\n"
"        f->b = mat->refl_b;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"        *specularBounce = mat->reflectionSpecularBounce;\n"
"#endif\n"
"    } else {\n"
"        const float kk = (into ? 1.f : -1.f) * (ddn * nnt + sqrt(cos2t));\n"
"        Vector nkk = *N;\n"
"        nkk.x *= kk;\n"
"        nkk.y *= kk;\n"
"        nkk.z *= kk;\n"
"\n"
"        Vector transDir;\n"
"        transDir.x = -nnt * wo->x - nkk.x;\n"
"        transDir.y = -nnt * wo->y - nkk.y;\n"
"        transDir.z = -nnt * wo->z - nkk.z;\n"
"        Normalize(&transDir);\n"
"\n"
"        const float c = 1.f - (into ? -ddn : Dot(&transDir, N));\n"
"\n"
"        const float R0 = mat->R0;\n"
"        const float Re = R0 + (1.f - R0) * c * c * c * c * c;\n"
"        const float Tr = 1.f - Re;\n"
"        const float P = .25f + .5f * Re;\n"
"\n"
"        if (Tr == 0.f) {\n"
"            if (Re == 0.f)\n"
"                *pdf = 0.f;\n"
"            else {\n"
"                *wi = reflDir;\n"
"                *pdf = 1.f;\n"
"\n"
"                f->r = mat->refl_r;\n"
"                f->g = mat->refl_g;\n"
"                f->b = mat->refl_b;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"                *specularBounce = mat->reflectionSpecularBounce;\n"
"#endif\n"
"            }\n"
"        } else if (Re == 0.f) {\n"
"            *wi = transDir;\n"
"            *pdf = 1.f;\n"
"\n"
"            f->r = mat->refrct_r;\n"
"            f->g = mat->refrct_g;\n"
"            f->b = mat->refrct_b;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            *specularBounce = mat->transmitionSpecularBounce;\n"
"#endif\n"
"        } else if (u0 < P) {\n"
"            *wi = reflDir;\n"
"            *pdf = P / Re;\n"
"\n"
"            f->r = mat->refl_r;\n"
"            f->g = mat->refl_g;\n"
"            f->b = mat->refl_b;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            *specularBounce = mat->reflectionSpecularBounce;\n"
"#endif\n"
"        } else {\n"
"            *wi = transDir;\n"
"            *pdf = (1.f - P) / Tr;\n"
"\n"
"            f->r = mat->refrct_r;\n"
"            f->g = mat->refrct_g;\n"
"            f->b = mat->refrct_b;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            *specularBounce = mat->transmitionSpecularBounce;\n"
"#endif\n"
"        }\n"
"    }\n"
"}\n"
"\n"
"void MatteMirror_Sample_f(__global MatteMirrorParam *mat, const Vector *wo, Vector *wi,\n"
"		float *pdf, Spectrum *f, const Vector *shadeN,\n"
"		const float u0, const float u1, const float u2\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"		) {\n"
"    const float totFilter = mat->totFilter;\n"
"    const float comp = u2 * totFilter;\n"
"\n"
"    if (comp > mat->matteFilter) {\n"
"        Mirror_Sample_f(&mat->mirror, wo, wi, pdf, f, shadeN\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            , specularBounce\n"
"#endif\n"
"            );\n"
"        *pdf *= mat->mirrorPdf;\n"
"    } else {\n"
"        Matte_Sample_f(&mat->matte, wo, wi, pdf, f, shadeN, u0, u1\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            , specularBounce\n"
"#endif\n"
"            );\n"
"        *pdf *= mat->mattePdf;\n"
"    }\n"
"}\n"
"\n"
"void GlossyReflection(const Vector *wo, Vector *wi, const float exponent,\n"
"		const Vector *shadeN, const float u0, const float u1) {\n"
"    const float phi = 2.f * M_PI * u0;\n"
"    const float cosTheta = pow(1.f - u1, exponent);\n"
"    const float sinTheta = sqrt(1.f - cosTheta * cosTheta);\n"
"    const float x = cos(phi) * sinTheta;\n"
"    const float y = sin(phi) * sinTheta;\n"
"    const float z = cosTheta;\n"
"\n"
"    Vector w;\n"
"    const float RdotShadeN = Dot(shadeN, wo);\n"
"	w.x = (2.f * RdotShadeN) * shadeN->x - wo->x;\n"
"	w.y = (2.f * RdotShadeN) * shadeN->y - wo->y;\n"
"	w.z = (2.f * RdotShadeN) * shadeN->z - wo->z;\n"
"\n"
"    Vector u, a;\n"
"    if (fabs(shadeN->x) > .1f) {\n"
"        a.x = 0.f;\n"
"        a.y = 1.f;\n"
"    } else {\n"
"        a.x = 1.f;\n"
"        a.y = 0.f;\n"
"    }\n"
"    a.z = 0.f;\n"
"    Cross(&u, &a, &w);\n"
"    Normalize(&u);\n"
"    Vector v;\n"
"    Cross(&v, &w, &u);\n"
"\n"
"    wi->x = x * u.x + y * v.x + z * w.x;\n"
"    wi->y = x * u.y + y * v.y + z * w.y;\n"
"    wi->z = x * u.z + y * v.z + z * w.z;\n"
"}\n"
"\n"
"void Metal_Sample_f(__global MetalParam *mat, const Vector *wo, Vector *wi,\n"
"		float *pdf, Spectrum *f, const Vector *shadeN,\n"
"		const float u0, const float u1\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"		) {\n"
"        GlossyReflection(wo, wi, mat->exponent, shadeN, u0, u1);\n"
"\n"
"		if (Dot(wi, shadeN) > 0.f) {\n"
"			*pdf = 1.f;\n"
"\n"
"            f->r = mat->r;\n"
"            f->g = mat->g;\n"
"            f->b = mat->b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            *specularBounce = mat->specularBounce;\n"
"#endif\n"
"		} else\n"
"			*pdf = 0.f;\n"
"}\n"
"\n"
"void MatteMetal_Sample_f(__global MatteMetalParam *mat, const Vector *wo, Vector *wi,\n"
"		float *pdf, Spectrum *f, const Vector *shadeN,\n"
"		const float u0, const float u1, const float u2\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"		) {\n"
"        const float totFilter = mat->totFilter;\n"
"        const float comp = u2 * totFilter;\n"
"\n"
"		if (comp > mat->matteFilter) {\n"
"            Metal_Sample_f(&mat->metal, wo, wi, pdf, f, shadeN, u0, u1\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"                , specularBounce\n"
"#endif\n"
"                );\n"
"			*pdf *= mat->metalPdf;\n"
"		} else {\n"
"            Matte_Sample_f(&mat->matte, wo, wi, pdf, f, shadeN, u0, u1\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"                , specularBounce\n"
"#endif\n"
"                );\n"
"			*pdf *= mat->mattePdf;\n"
"		}\n"
"}\n"
"\n"
"void Alloy_Sample_f(__global AlloyParam *mat, const Vector *wo, Vector *wi,\n"
"		float *pdf, Spectrum *f, const Vector *shadeN,\n"
"		const float u0, const float u1, const float u2\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"		) {\n"
"    // Schilick's approximation\n"
"    const float c = 1.f - Dot(wo, shadeN);\n"
"    const float R0 = mat->R0;\n"
"    const float Re = R0 + (1.f - R0) * c * c * c * c * c;\n"
"\n"
"    const float P = .25f + .5f * Re;\n"
"\n"
"    if (u2 < P) {\n"
"        GlossyReflection(wo, wi, mat->exponent, shadeN, u0, u1);\n"
"        *pdf = P / Re;\n"
"\n"
"        f->r = Re * mat->refl_r;\n"
"        f->g = Re * mat->refl_g;\n"
"        f->b = Re * mat->refl_b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"        *specularBounce = mat->specularBounce;\n"
"#endif\n"
"    } else {\n"
"        Vector dir;\n"
"        CosineSampleHemisphere(&dir, u0, u1);\n"
"        *pdf = dir.z * INV_PI;\n"
"\n"
"        Vector v1, v2;\n"
"        CoordinateSystem(shadeN, &v1, &v2);\n"
"\n"
"        wi->x = v1.x * dir.x + v2.x * dir.y + shadeN->x * dir.z;\n"
"        wi->y = v1.y * dir.x + v2.y * dir.y + shadeN->y * dir.z;\n"
"        wi->z = v1.z * dir.x + v2.z * dir.y + shadeN->z * dir.z;\n"
"\n"
"        const float dp = Dot(shadeN, wi);\n"
"        // Using 0.0001 instead of 0.0 to cut down fireflies\n"
"        if (dp <= 0.0001f)\n"
"            *pdf = 0.f;\n"
"        else {\n"
"            *pdf /=  dp;\n"
"\n"
"            const float iRe = 1.f - Re;\n"
"            *pdf *= (1.f - P) / iRe;\n"
"\n"
"            f->r = iRe * mat->diff_r;\n"
"            f->g = iRe * mat->diff_g;\n"
"            f->b = iRe * mat->diff_b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            *specularBounce = FALSE;\n"
"#endif\n"
"        }\n"
"    }\n"
"}\n"
"\n"
"void ArchGlass_Sample_f(__global ArchGlassParam *mat,\n"
"    const Vector *wo, Vector *wi, float *pdf, Spectrum *f, const Vector *N, const Vector *shadeN,\n"
"    const float u0\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global int *specularBounce\n"
"#endif\n"
"        ) {\n"
"    // Ray from outside going in ?\n"
"    const bool into = (Dot(N, shadeN) > 0.f);\n"
"\n"
"    if (!into) {\n"
"        // No internal reflections\n"
"        wi->x = -wo->x;\n"
"        wi->y = -wo->y;\n"
"        wi->z = -wo->z;\n"
"        *pdf = 1.f;\n"
"\n"
"        f->r = mat->refrct_r;\n"
"        f->g = mat->refrct_g;\n"
"        f->b = mat->refrct_b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"        *specularBounce = mat->transmitionSpecularBounce;\n"
"#endif\n"
"    } else {\n"
"        // RR to choose if reflect the ray or go trough the glass\n"
"        const float comp = u0 * mat->totFilter;\n"
"\n"
"        if (comp > mat->transFilter) {\n"
"            const float k = 2.f * Dot(N, wo);\n"
"            wi->x = k * N->x - wo->x;\n"
"            wi->y = k * N->y - wo->y;\n"
"            wi->z = k * N->z - wo->z;\n"
"            *pdf =  mat->reflPdf;\n"
"\n"
"            f->r = mat->refl_r;\n"
"            f->g = mat->refl_g;\n"
"            f->b = mat->refl_b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            *specularBounce = mat->reflectionSpecularBounce;\n"
"#endif\n"
"        } else {\n"
"            wi->x = -wo->x;\n"
"            wi->y = -wo->y;\n"
"            wi->z = -wo->z;\n"
"            *pdf =  mat->transPdf;\n"
"\n"
"            f->r = mat->refrct_r;\n"
"            f->g = mat->refrct_g;\n"
"            f->b = mat->refrct_b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"            *specularBounce = mat->transmitionSpecularBounce;\n"
"#endif\n"
"        }\n"
"    }\n"
"}\n"
"\n"
"//------------------------------------------------------------------------------\n"
"// Lights\n"
"//------------------------------------------------------------------------------\n"
"\n"
"void AreaLight_Le(__global AreaLightParam *mat, const Vector *wo, const Vector *lightN, Spectrum *Le) {\n"
"	if (Dot(lightN, wo) > 0.f) {\n"
"        Le->r = mat->gain_r;\n"
"        Le->g = mat->gain_g;\n"
"        Le->b = mat->gain_b;\n"
"    }\n"
"}\n"
"\n"
"void SampleTriangleLight(__global TriangleLight *light,	const float u0, const float u1, Point *p) {\n"
"	Point v0, v1, v2;\n"
"	v0 = light->v0;\n"
"	v1 = light->v1;\n"
"	v2 = light->v2;\n"
"\n"
"	// UniformSampleTriangle(u0, u1, b0, b1);\n"
"	const float su1 = sqrt(u0);\n"
"	const float b0 = 1.f - su1;\n"
"	const float b1 = u1 * su1;\n"
"	const float b2 = 1.f - b0 - b1;\n"
"\n"
"	p->x = b0 * v0.x + b1 * v1.x + b2 * v2.x;\n"
"	p->y = b0 * v0.y + b1 * v1.y + b2 * v2.y;\n"
"	p->z = b0 * v0.z + b1 * v1.z + b2 * v2.z;\n"
"}\n"
"\n"
"void TriangleLight_Sample_L(__global TriangleLight *l,\n"
"		const Vector *wo, const Point *hitPoint,\n"
"		float *pdf, Spectrum *f, Ray *shadowRay,\n"
"		const float u0, const float u1, const float u2) {\n"
"	Point samplePoint;\n"
"	SampleTriangleLight(l, u0, u1, &samplePoint);\n"
"\n"
"	shadowRay->d.x = samplePoint.x - hitPoint->x;\n"
"	shadowRay->d.y = samplePoint.y - hitPoint->y;\n"
"	shadowRay->d.z = samplePoint.z - hitPoint->z;\n"
"	const float distanceSquared = Dot(&shadowRay->d, &shadowRay->d);\n"
"	const float distance = sqrt(distanceSquared);\n"
"	const float invDistance = 1.f / distance;\n"
"	shadowRay->d.x *= invDistance;\n"
"	shadowRay->d.y *= invDistance;\n"
"	shadowRay->d.z *= invDistance;\n"
"\n"
"	Vector sampleN = l->normal;\n"
"	const float sampleNdotMinusWi = -Dot(&sampleN, &shadowRay->d);\n"
"	if (sampleNdotMinusWi <= 0.f)\n"
"		*pdf = 0.f;\n"
"	else {\n"
"		*pdf = distanceSquared / (sampleNdotMinusWi * l->area);\n"
"\n"
"		// Using 0.1 instead of 0.0 to cut down fireflies\n"
"		if (*pdf <= 0.1f)\n"
"			*pdf = 0.f;\n"
"		else {\n"
"            shadowRay->o = *hitPoint;\n"
"            shadowRay->mint = PARAM_RAY_EPSILON;\n"
"            shadowRay->maxt = distance - PARAM_RAY_EPSILON;\n"
"\n"
"            f->r = l->gain_r;\n"
"            f->g = l->gain_g;\n"
"            f->b = l->gain_b;\n"
"        }\n"
"	}\n"
"}\n"
"\n"
"//------------------------------------------------------------------------------\n"
"\n"
"void TerminatePath(__global Path *path, __global Ray *ray, __global Pixel *frameBuffer, Seed *seed, Spectrum *radiance\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"		, __global float *cameraData\n"
"#endif\n"
"		) {\n"
"	// Add sample to the framebuffer\n"
"\n"
"	const uint pixelIndex = path->pixelIndex;\n"
"	__global Pixel *pixel = &frameBuffer[pixelIndex];\n"
"\n"
"	pixel->c.r += isnan(radiance->r) ? 0.f : radiance->r;\n"
"	pixel->c.g += isnan(radiance->g) ? 0.f : radiance->g;\n"
"	pixel->c.b += isnan(radiance->b) ? 0.f : radiance->b;\n"
"	pixel->count += 1;\n"
"\n"
"    // Re-initialize the path\n"
"\n"
"	uint newPixelIndex;\n"
"#if (PARAM_SAMPLE_PER_PIXEL > 0)\n"
"	newPixelIndex = (pixelIndex + PARAM_PATH_COUNT) % (PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT);\n"
"	path->pixelIndex = newPixelIndex;\n"
"#else\n"
"	const uint subpixelIndex = path->subpixelIndex;\n"
"	if (subpixelIndex >= PARAM_SAMPLE_PER_PIXEL) {\n"
"		newPixelIndex = (pixelIndex + PARAM_PATH_COUNT) % (PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT);\n"
"		path->pixelIndex = newPixelIndex;\n"
"		path->subpixelIndex = 0;\n"
"	} else {\n"
"		newPixelIndex = pixelIndex;\n"
"		path->subpixelIndex = subpixelIndex + 1;\n"
"	}\n"
"#endif\n"
"\n"
"	GenerateRay(newPixelIndex, ray, seed\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"		, cameraData\n"
"#endif\n"
"		);\n"
"\n"
"	path->throughput.r = 1.f;\n"
"	path->throughput.g = 1.f;\n"
"	path->throughput.b = 1.f;\n"
"	path->depth = 0;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"	path->specularBounce = TRUE;\n"
"	path->state = PATH_STATE_NEXT_VERTEX;\n"
"	path->accumRadiance.r = 0.f;\n"
"	path->accumRadiance.g = 0.f;\n"
"	path->accumRadiance.b = 0.f;\n"
"#endif\n"
"}\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"__kernel void AdvancePaths_Step1(\n"
"		__global Path *paths,\n"
"		__global Ray *rays,\n"
"		__global RayHit *rayHits,\n"
"		__global Material *mats,\n"
"		__global uint *meshMats,\n"
"		__global uint *meshIDs,\n"
"		__global Spectrum *vertColors,\n"
"		__global Vector *vertNormals,\n"
"		__global Triangle *triangles,\n"
"        __global TriangleLight *triLights\n"
"#if defined(PARAM_HAS_TEXTUREMAPS)\n"
"        , __global Spectrum *texMapRGBBuff\n"
"#if defined(PARAM_HAS_ALPHA_TEXTUREMAPS)\n"
"		, __global float *texMapAlphaBuff\n"
"#endif\n"
"        , __global TexMap *texMapDescBuff\n"
"        , __global unsigned int *meshTexsBuff\n"
"        , __global UV *vertUVs\n"
"#endif\n"
"		) {\n"
"	const int gid = get_global_id(0);\n"
"	__global Path *path = &paths[gid];\n"
"\n"
"	__global Ray *ray = &rays[gid];\n"
"	__global RayHit *rayHit = &rayHits[gid];\n"
"	uint currentTriangleIndex = rayHit->index;\n"
"\n"
"	// Read the seed\n"
"	Seed seed;\n"
"	seed.s1 = path->seed.s1;\n"
"	seed.s2 = path->seed.s2;\n"
"	seed.s3 = path->seed.s3;\n"
"\n"
"	if (path->state == PATH_STATE_SAMPLE_LIGHT) {\n"
"		if (currentTriangleIndex == 0xffffffffu) {\n"
"			// Nothing was hit, the light is visible\n"
"\n"
"			path->accumRadiance.r += path->lightRadiance.r;\n"
"			path->accumRadiance.g += path->lightRadiance.g;\n"
"			path->accumRadiance.b += path->lightRadiance.b;\n"
"		}\n"
"#if defined(PARAM_HAS_ALPHA_TEXTUREMAPS) || defined(PARAM_ENABLE_MAT_ARCHGLASS)\n"
"		else {\n"
"			const uint meshIndex = meshIDs[currentTriangleIndex];\n"
"\n"
"#if defined(PARAM_HAS_ALPHA_TEXTUREMAPS)\n"
"			// Check if is a mesh with alpha texture map applied\n"
"\n"
"			const float hitPointB1 = rayHit->b1;\n"
"			const float hitPointB2 = rayHit->b2;\n"
"\n"
"			// Check it the mesh has a texture map\n"
"			unsigned int texIndex = meshTexsBuff[meshIndex];\n"
"			if (texIndex != 0xffffffffu) {\n"
"				__global TexMap *texMap = &texMapDescBuff[texIndex];\n"
"\n"
"				// Check if it has an alpha channel\n"
"				if (texMap->alphaOffset != 0xffffffffu) {\n"
"					// Interpolate UV coordinates\n"
"					UV uv;\n"
"					Mesh_InterpolateUV(vertUVs, triangles, currentTriangleIndex, hitPointB1, hitPointB2, &uv);\n"
"\n"
"					const float alpha = TexMap_GetAlpha(&texMapAlphaBuff[texMap->alphaOffset], texMap->width, texMap->height, uv.u, uv.v);\n"
"\n"
"					if ((alpha == 0.0f) || ((alpha < 1.f) && (RndFloatValue(&seed) > alpha))) {\n"
"						// Continue to trace the ray\n"
"						const float hitPointT = rayHit->t;\n"
"						ray->o.x = ray->o.x + ray->d.x * hitPointT;\n"
"						ray->o.y = ray->o.y + ray->d.y * hitPointT;\n"
"						ray->o.z = ray->o.z + ray->d.z * hitPointT;\n"
"						ray->maxt -= hitPointT;\n"
"\n"
"						// Save the seed\n"
"						path->seed.s1 = seed.s1;\n"
"						path->seed.s2 = seed.s2;\n"
"						path->seed.s3 = seed.s3;\n"
"\n"
"						return;\n"
"					}\n"
"				}\n"
"			}\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_ARCHGLASS)\n"
"            // Check if is a shadow transparent material\n"
"            __global Material *hitPointMat = &mats[meshMats[meshIndex]];\n"
"\n"
"            if (hitPointMat->type == MAT_ARCHGLASS) {\n"
"                const float hitPointT = rayHit->t;\n"
"                ray->o.x = ray->o.x + ray->d.x * hitPointT;\n"
"                ray->o.y = ray->o.y + ray->d.y * hitPointT;\n"
"                ray->o.z = ray->o.z + ray->d.z * hitPointT;\n"
"                ray->maxt -= hitPointT;\n"
"\n"
"                path->lightRadiance.r *= hitPointMat->param.archGlass.refrct_r;\n"
"                path->lightRadiance.g *= hitPointMat->param.archGlass.refrct_g;\n"
"                path->lightRadiance.b *= hitPointMat->param.archGlass.refrct_b;\n"
"\n"
"                // Continue to trace the ray\n"
"                return;\n"
"            }\n"
"#endif\n"
"        }\n"
"#endif\n"
"\n"
"		// Restore the path RayHit\n"
"		*rayHit = path->pathHit;\n"
"\n"
"		// Restore the path Ray\n"
"        *ray = path->pathRay;\n"
"\n"
"        path->state = PATH_STATE_NEXT_VERTEX;\n"
"        return;\n"
"	}\n"
"\n"
"    // state is PATH_STATE_NEXT_VERTEX or PATH_STATE_CONTINUE_NEXT_VERTEX\n"
"\n"
"	// If we are in PATH_STATE_CONTINUE_NEXT_VERTEX state\n"
"	path->state = PATH_STATE_NEXT_VERTEX;\n"
"\n"
"    if (currentTriangleIndex != 0xffffffffu) {\n"
"        // Something was hit\n"
"\n"
"		const uint meshIndex = meshIDs[currentTriangleIndex];\n"
"		__global Material *hitPointMat = &mats[meshMats[meshIndex]];\n"
"\n"
"        Vector rayDir = ray->d;\n"
"\n"
"        const float hitPointT = rayHit->t;\n"
"        const float hitPointB1 = rayHit->b1;\n"
"        const float hitPointB2 = rayHit->b2;\n"
"\n"
"        Point hitPoint;\n"
"        hitPoint.x = ray->o.x + rayDir.x * hitPointT;\n"
"        hitPoint.y = ray->o.y + rayDir.y * hitPointT;\n"
"        hitPoint.z = ray->o.z + rayDir.z * hitPointT;\n"
"\n"
"		// Interpolate Color\n"
"		Spectrum shadeColor;\n"
"		Mesh_InterpolateColor(vertColors, triangles, currentTriangleIndex, hitPointB1, hitPointB2, &shadeColor);\n"
"\n"
"#if defined(PARAM_HAS_TEXTUREMAPS)\n"
"        // Interpolate UV coordinates\n"
"        UV uv;\n"
"        Mesh_InterpolateUV(vertUVs, triangles, currentTriangleIndex, hitPointB1, hitPointB2, &uv);\n"
"\n"
"        // Check it the mesh has a texture map\n"
"        unsigned int texIndex = meshTexsBuff[meshIndex];\n"
"        if (texIndex != 0xffffffffu) {\n"
"            __global TexMap *texMap = &texMapDescBuff[texIndex];\n"
"\n"
"#if defined(PARAM_HAS_ALPHA_TEXTUREMAPS)\n"
"			// Check if it has an alpha channel\n"
"			if (texMap->alphaOffset != 0xffffffffu) {\n"
"				const float alpha = TexMap_GetAlpha(&texMapAlphaBuff[texMap->alphaOffset], texMap->width, texMap->height, uv.u, uv.v);\n"
"\n"
"				if ((alpha == 0.0f) || ((alpha < 1.f) && (RndFloatValue(&seed) > alpha))) {\n"
"					// Continue to trace the ray\n"
"					ray->o = hitPoint;\n"
"					ray->maxt -= hitPointT;\n"
"\n"
"					path->state = PATH_STATE_CONTINUE_NEXT_VERTEX;\n"
"					// Save the seed\n"
"					path->seed.s1 = seed.s1;\n"
"					path->seed.s2 = seed.s2;\n"
"					path->seed.s3 = seed.s3;\n"
"\n"
"					return;\n"
"				}\n"
"			}\n"
"#endif\n"
"\n"
"            Spectrum texColor;\n"
"            TexMap_GetColor(&texMapRGBBuff[texMap->rgbOffset], texMap->width, texMap->height, uv.u, uv.v, &texColor);\n"
"\n"
"            shadeColor.r *= texColor.r;\n"
"            shadeColor.g *= texColor.g;\n"
"            shadeColor.b *= texColor.b;\n"
"        }\n"
"#endif\n"
"\n"
"		// Interpolate the normal\n"
"        Vector N;\n"
"		Mesh_InterpolateNormal(vertNormals, triangles, currentTriangleIndex, hitPointB1, hitPointB2, &N);\n"
"\n"
"		// Flip the normal if required\n"
"        Vector shadeN;\n"
"		const float nFlip = (Dot(&rayDir, &N) > 0.f) ? -1.f : 1.f;\n"
"		shadeN.x = nFlip * N.x;\n"
"		shadeN.y = nFlip * N.y;\n"
"		shadeN.z = nFlip * N.z;\n"
"\n"
"        float directLightPdf;\n"
"        switch (hitPointMat->type) {\n"
"            case MAT_MATTE:\n"
"                directLightPdf = 1.f;\n"
"                break;\n"
"            case MAT_MATTEMIRROR:\n"
"                directLightPdf = 1.f / hitPointMat->param.matteMirror.mattePdf;\n"
"                break;\n"
"            case MAT_MATTEMETAL:\n"
"                directLightPdf = 1.f / hitPointMat->param.matteMetal.mattePdf;\n"
"                break;\n"
"            case MAT_ALLOY: {\n"
"                // Schilick's approximation\n"
"                const float c = 1.f + Dot(&rayDir, &shadeN);\n"
"                const float R0 = hitPointMat->param.alloy.R0;\n"
"                const float Re = R0 + (1.f - R0) * c * c * c * c * c;\n"
"\n"
"                const float P = .25f + .5f * Re;\n"
"\n"
"                directLightPdf = (1.f - P) / (1.f - Re);\n"
"                break;\n"
"            }\n"
"            default:\n"
"                directLightPdf = 0.f;\n"
"                break;\n"
"        }\n"
"\n"
"        if (directLightPdf > 0.f) {\n"
"            // Select a light source to sample\n"
"            const uint lightIndex = min((uint)floor(PARAM_DL_LIGHT_COUNT * RndFloatValue(&seed)), (uint)(PARAM_DL_LIGHT_COUNT - 1));\n"
"            __global TriangleLight *l = &triLights[lightIndex];\n"
"\n"
"            // Setup the shadow ray\n"
"            Vector wo;\n"
"            wo.x = -rayDir.x;\n"
"            wo.y = -rayDir.y;\n"
"            wo.z = -rayDir.z;\n"
"            Spectrum Le;\n"
"            float lightPdf;\n"
"            Ray shadowRay;\n"
"            TriangleLight_Sample_L(l, &wo, &hitPoint, &lightPdf, &Le, &shadowRay,\n"
"                RndFloatValue(&seed), RndFloatValue(&seed), RndFloatValue(&seed));\n"
"\n"
"            const float dp = Dot(&shadeN, &shadowRay.d);\n"
"            const float matPdf = (dp <= 0.f) ? 0.f : 1.f;\n"
"\n"
"            const float pdf = lightPdf * matPdf * directLightPdf;\n"
"            if (pdf > 0.f) {\n"
"                Spectrum throughput = path->throughput;\n"
"                throughput.r *= shadeColor.r;\n"
"                throughput.g *= shadeColor.g;\n"
"                throughput.b *= shadeColor.b;\n"
"\n"
"                const float k = dp * PARAM_DL_LIGHT_COUNT / (pdf * M_PI);\n"
"                // NOTE: I assume all matte mixed material have a MatteParam as first field\n"
"                path->lightRadiance.r = throughput.r * hitPointMat->param.matte.r * k * Le.r;\n"
"                path->lightRadiance.g = throughput.g * hitPointMat->param.matte.g * k * Le.g;\n"
"                path->lightRadiance.b = throughput.b * hitPointMat->param.matte.b * k * Le.b;\n"
"\n"
"                // Save current ray hit information\n"
"                path->pathHit.t = hitPointT;\n"
"                path->pathHit.b1 = hitPointB1;\n"
"                path->pathHit.b2 = hitPointB2;\n"
"                path->pathHit.index = currentTriangleIndex;\n"
"\n"
"                // Save the current Ray\n"
"                path->pathRay = *ray;\n"
"\n"
"                *ray = shadowRay;\n"
"\n"
"                path->state = PATH_STATE_SAMPLE_LIGHT;\n"
"            }\n"
"\n"
"            // Save the seed\n"
"            path->seed.s1 = seed.s1;\n"
"            path->seed.s2 = seed.s2;\n"
"            path->seed.s3 = seed.s3;\n"
"        }\n"
"    }\n"
"}\n"
"#endif\n"
"\n"
"__kernel void AdvancePaths_Step2(\n"
"		__global Path *paths,\n"
"		__global Ray *rays,\n"
"		__global RayHit *rayHits,\n"
"		__global Pixel *frameBuffer,\n"
"		__global Material *mats,\n"
"		__global uint *meshMats,\n"
"		__global uint *meshIDs,\n"
"		__global Spectrum *vertColors,\n"
"		__global Vector *vertNormals,\n"
"		__global Triangle *triangles\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"		, __global float *cameraData\n"
"#endif\n"
"#if defined(PARAM_HAVE_INFINITELIGHT)\n"
"		, __global Spectrum *infiniteLightMap\n"
"#endif\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		, __global TriangleLight *triLights\n"
"#endif\n"
"#if defined(PARAM_HAS_TEXTUREMAPS)\n"
"        , __global Spectrum *texMapRGBBuff\n"
"#if defined(PARAM_HAS_ALPHA_TEXTUREMAPS)\n"
"		, __global float *texMapAlphaBuff\n"
"#endif\n"
"        , __global TexMap *texMapDescBuff\n"
"        , __global unsigned int *meshTexsBuff\n"
"        , __global UV *vertUVs\n"
"#endif\n"
"		) {\n"
"	const int gid = get_global_id(0);\n"
"	__global Path *path = &paths[gid];\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"	int pathState = path->state;\n"
"	if ((pathState == PATH_STATE_SAMPLE_LIGHT) || (pathState == PATH_STATE_CONTINUE_NEXT_VERTEX)) {\n"
"        // Handled by Step1 kernel\n"
"        return;\n"
"    }\n"
"#endif\n"
"\n"
"	// Read the seed\n"
"	Seed seed;\n"
"	seed.s1 = path->seed.s1;\n"
"	seed.s2 = path->seed.s2;\n"
"	seed.s3 = path->seed.s3;\n"
"\n"
"	__global Ray *ray = &rays[gid];\n"
"	__global RayHit *rayHit = &rayHits[gid];\n"
"	uint currentTriangleIndex = rayHit->index;\n"
"\n"
"	const float hitPointT = rayHit->t;\n"
"    const float hitPointB1 = rayHit->b1;\n"
"    const float hitPointB2 = rayHit->b2;\n"
"\n"
"    Vector rayDir = ray->d;\n"
"\n"
"	Point hitPoint;\n"
"    hitPoint.x = ray->o.x + rayDir.x * hitPointT;\n"
"    hitPoint.y = ray->o.y + rayDir.y * hitPointT;\n"
"    hitPoint.z = ray->o.z + rayDir.z * hitPointT;\n"
"\n"
"	Spectrum throughput = path->throughput;\n"
"\n"
"    if (currentTriangleIndex != 0xffffffffu) {\n"
"		// Something was hit\n"
"\n"
"		const uint meshIndex = meshIDs[currentTriangleIndex];\n"
"		__global Material *hitPointMat = &mats[meshMats[meshIndex]];\n"
"		uint matType = hitPointMat->type;\n"
"\n"
"		// Interpolate Color\n"
"        Spectrum shadeColor;\n"
"		Mesh_InterpolateColor(vertColors, triangles, currentTriangleIndex, hitPointB1, hitPointB2, &shadeColor);\n"
"\n"
"#if defined(PARAM_HAS_TEXTUREMAPS)\n"
"        // Interpolate UV coordinates\n"
"        UV uv;\n"
"        Mesh_InterpolateUV(vertUVs, triangles, currentTriangleIndex, hitPointB1, hitPointB2, &uv);\n"
"\n"
"        // Check it the mesh has a texture map\n"
"        unsigned int texIndex = meshTexsBuff[meshIndex];\n"
"        if (texIndex != 0xffffffffu) {\n"
"            __global TexMap *texMap = &texMapDescBuff[texIndex];\n"
"\n"
"#if defined(PARAM_HAS_ALPHA_TEXTUREMAPS) && !defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"			// Check if it has an alpha channel\n"
"			if (texMap->alphaOffset != 0xffffffffu) {\n"
"				const float alpha = TexMap_GetAlpha(&texMapAlphaBuff[texMap->alphaOffset], texMap->width, texMap->height, uv.u, uv.v);\n"
"\n"
"				if ((alpha == 0.0f) || ((alpha < 1.f) && (RndFloatValue(&seed) > alpha))) {\n"
"					// Continue to trace the ray\n"
"					matType = MAT_NULL;\n"
"				}\n"
"			}\n"
"#endif\n"
"\n"
"            Spectrum texColor;\n"
"            TexMap_GetColor(&texMapRGBBuff[texMap->rgbOffset], texMap->width, texMap->height, uv.u, uv.v, &texColor);\n"
"\n"
"            shadeColor.r *= texColor.r;\n"
"            shadeColor.g *= texColor.g;\n"
"            shadeColor.b *= texColor.b;\n"
"        }\n"
"#endif\n"
"\n"
"		throughput.r *= shadeColor.r;\n"
"		throughput.g *= shadeColor.g;\n"
"		throughput.b *= shadeColor.b;\n"
"\n"
"		// Interpolate the normal\n"
"        Vector N;\n"
"		Mesh_InterpolateNormal(vertNormals, triangles, currentTriangleIndex, hitPointB1, hitPointB2, &N);\n"
"\n"
"		// Flip the normal if required\n"
"        Vector shadeN;\n"
"		const float nFlip = (Dot(&rayDir, &N) > 0.f) ? -1.f : 1.f;\n"
"		shadeN.x = nFlip * N.x;\n"
"		shadeN.y = nFlip * N.y;\n"
"		shadeN.z = nFlip * N.z;\n"
"\n"
"		const float u0 = RndFloatValue(&seed);\n"
"		const float u1 = RndFloatValue(&seed);\n"
"		const float u2 = RndFloatValue(&seed);\n"
"\n"
"		Vector wo;\n"
"		wo.x = -rayDir.x;\n"
"		wo.y = -rayDir.y;\n"
"		wo.z = -rayDir.z;\n"
"\n"
"		Vector wi;\n"
"		float pdf;\n"
"		Spectrum f;\n"
"\n"
"		switch (matType) {\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_MATTE)\n"
"			case MAT_MATTE:\n"
"				Matte_Sample_f(&hitPointMat->param.matte, &wo, &wi, &pdf, &f, &shadeN, u0, u1\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"				break;\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_AREALIGHT)\n"
"			case MAT_AREALIGHT: {\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"				if (path->specularBounce) {\n"
"#endif\n"
"					Spectrum radiance;\n"
"					AreaLight_Le(&hitPointMat->param.areaLight, &wo, &N, &radiance);\n"
"                    radiance.r *= throughput.r;\n"
"                    radiance.g *= throughput.g;\n"
"                    radiance.b *= throughput.b;\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					radiance.r += path->accumRadiance.r;\n"
"					radiance.g += path->accumRadiance.g;\n"
"					radiance.b += path->accumRadiance.b;\n"
"#endif\n"
"\n"
"					TerminatePath(path, ray, frameBuffer, &seed, &radiance\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"						, cameraData\n"
"#endif\n"
"					);\n"
"\n"
"					// Save the seed\n"
"					path->seed.s1 = seed.s1;\n"
"					path->seed.s2 = seed.s2;\n"
"					path->seed.s3 = seed.s3;\n"
"\n"
"					return;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"				}\n"
"				break;\n"
"#endif\n"
"			}\n"
"#endif\n"
"\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_MIRROR)\n"
"			case MAT_MIRROR:\n"
"				Mirror_Sample_f(&hitPointMat->param.mirror, &wo, &wi, &pdf, &f, &shadeN\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"				break;\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_GLASS)\n"
"			case MAT_GLASS:\n"
"				Glass_Sample_f(&hitPointMat->param.glass, &wo, &wi, &pdf, &f, &N, &shadeN, u0\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"				break;\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_MATTEMIRROR)\n"
"            case MAT_MATTEMIRROR:\n"
"                MatteMirror_Sample_f(&hitPointMat->param.matteMirror, &wo, &wi, &pdf, &f, &shadeN, u0, u1, u2\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"                break;\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_METAL)\n"
"            case MAT_METAL:\n"
"				Metal_Sample_f(&hitPointMat->param.metal, &wo, &wi, &pdf, &f, &shadeN, u0, u1\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"                break;\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_MATTEMETAL)\n"
"            case MAT_MATTEMETAL:\n"
"                MatteMetal_Sample_f(&hitPointMat->param.matteMetal, &wo, &wi, &pdf, &f, &shadeN, u0, u1, u2\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"                break;\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_ALLOY)\n"
"            case MAT_ALLOY:\n"
"				Alloy_Sample_f(&hitPointMat->param.alloy, &wo, &wi, &pdf, &f, &shadeN, u0, u1, u2\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"                break;\n"
"#endif\n"
"\n"
"#if defined(PARAM_ENABLE_MAT_ARCHGLASS)\n"
"            case MAT_ARCHGLASS:\n"
"				ArchGlass_Sample_f(&hitPointMat->param.archGlass, &wo, &wi, &pdf, &f, &N, &shadeN, u0\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"					, &path->specularBounce\n"
"#endif\n"
"					);\n"
"                break;\n"
"#endif\n"
"\n"
"			case MAT_NULL:\n"
"				wi = rayDir;\n"
"				pdf = 1.f;\n"
"				f.r = 1.f;\n"
"				f.g = 1.f;\n"
"				f.b = 1.f;\n"
"\n"
"				// I have also to restore the original throughput\n"
"				throughput = path->throughput;\n"
"				break;\n"
"\n"
"			default:\n"
"				// Huston, we have a problem...\n"
"				pdf = 0.f;\n"
"				break;\n"
"		}\n"
"\n"
"		const uint pathDepth = path->depth + 1;\n"
"        const float invPdf = ((pdf <= 0.f) || (pathDepth >= PARAM_MAX_PATH_DEPTH)) ? 0.f : (1.f / pdf);\n"
"        throughput.r *= f.r * invPdf;\n"
"		throughput.g *= f.g * invPdf;\n"
"		throughput.b *= f.b * invPdf;\n"
"\n"
"        // Russian roulette\n"
"        const float rrProb = max(max(throughput.r, max(throughput.g, throughput.b)), (float)PARAM_RR_CAP);\n"
"        const float rrSample = RndFloatValue(&seed);\n"
"        const float invRRProb = (pathDepth > PARAM_RR_DEPTH) ? ((rrProb >= rrSample) ? 0.f : (1.f / rrProb)) : 1.f;\n"
"        throughput.r *= invRRProb;\n"
"        throughput.g *= invRRProb;\n"
"        throughput.b *= invRRProb;\n"
"\n"
"		if ((throughput.r <= 0.f) && (throughput.g <= 0.f) && (throughput.b <= 0.f)) {\n"
"			Spectrum radiance;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"			radiance.r = path->accumRadiance.r;\n"
"			radiance.g = path->accumRadiance.g;\n"
"			radiance.b = path->accumRadiance.b;\n"
"#else\n"
"			radiance.r = 0.f;\n"
"			radiance.g = 0.f;\n"
"			radiance.b = 0.f;\n"
"#endif\n"
"\n"
"			TerminatePath(path, ray, frameBuffer, &seed, &radiance\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"				, cameraData\n"
"#endif\n"
"				);\n"
"		} else {\n"
"			path->throughput = throughput;\n"
"\n"
"			// Setup next ray\n"
"			ray->o = hitPoint;\n"
"			ray->d = wi;\n"
"			ray->mint = PARAM_RAY_EPSILON;\n"
"			ray->maxt = FLT_MAX;\n"
"\n"
"			path->depth = pathDepth;\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"			path->state = PATH_STATE_NEXT_VERTEX;\n"
"#endif\n"
"		}\n"
"	} else {\n"
"		Spectrum radiance;\n"
"\n"
"#if defined(PARAM_HAVE_INFINITELIGHT)\n"
"		Spectrum Le;\n"
"		InfiniteLight_Le(infiniteLightMap, &Le, &rayDir);\n"
"\n"
"		radiance.r = throughput.r * Le.r;\n"
"		radiance.g = throughput.g * Le.g;\n"
"		radiance.b = throughput.b * Le.b;\n"
"#else\n"
"		radiance.r = 0.f;\n"
"		radiance.g = 0.f;\n"
"		radiance.b = 0.f;\n"
"#endif\n"
"\n"
"#if defined(PARAM_DIRECT_LIGHT_SAMPLING)\n"
"		radiance.r += path->accumRadiance.r;\n"
"		radiance.g += path->accumRadiance.g;\n"
"		radiance.b += path->accumRadiance.b;\n"
"#endif\n"
"\n"
"		TerminatePath(path, ray, frameBuffer, &seed, &radiance\n"
"#if defined(PARAM_CAMERA_DYNAMIC)\n"
"				, cameraData\n"
"#endif\n"
"			);\n"
"	}\n"
"\n"
"	// Save the seed\n"
"	path->seed.s1 = seed.s1;\n"
"	path->seed.s2 = seed.s2;\n"
"	path->seed.s3 = seed.s3;\n"
"}\n"
;
